{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<div align=\"center\"><h1>CUDA C/C++ を使用したアプリケーションの高速化</h1></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![CUDA](./images/CUDA_Logo.jpg)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "CPU のみのコンピューティングに置き換わりつつあるベスト プラクティスが、アクセラレーテッド コンピューティングです。さまざまな技術の飛躍的な進歩、高まるアクセラレーテッド アプリケーション需要、アプリケーション開発を容易にするプログラミング規則、アプリケーションをサポートするハードウェアの継続的な改良などにより、この新しいコンピューティング手法は急速に普及しつつあります。\n",
    "\n",
    "アクセラレーテッド コンピューティングの人気の理由が、驚異的なパフォーマンスと使いやすさを備えた  [CUDA](https://developer.nvidia.com/about-cuda) コンピューティング プラットフォームです。C、C++、Python、Fortran などの言語を拡張するコーディング パラダイムにより、世界最高の性能を誇る並列プロセッサである NVIDIA GPU 上でアクセラレートされた大規模並列化コードを実行できます。CUDA は、[DNN](https://developer.nvidia.com/cudnn)、[BLAS](https://developer.nvidia.com/cublas)、[グラフ分析](https://developer.nvidia.com/nvgraph)、[FFT](https://developer.nvidia.com/cufft) などの高度に最適化されたライブラリのエコシステムに加え、強力な [コマンド ライン](http://docs.nvidia.com/cuda/profiler-users-guide/index.html#nvprof-overview)と[ビジュアル プロファイラー](http://docs.nvidia.com/cuda/profiler-users-guide/index.html#visual)が付属しており、手間をかけずにアプリケーションを大幅に高速化できます。\n",
    "\n",
    "CUDA は、[計算流体力学](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=10,12,16,17,19,51,53,71,87,121,124,156,157,195,202,203,204,312,339,340,395,407,448,485,517,528,529,541,245,216,104,462,513,250,492,420,429,490,10,12,16,17,19,51,53,71,87,121,124,156,157,195,202,203,204,312,339,340,395,407,448,485,517,528,529,541,245,216,104,462,513,250,492,420,429,490,10,12,16,17,19,51,53,71,87,121,124,156,157,195,202,203,204,312,339,340,395,407,448,485,517,528,529,541,245,216,104,462,513,250,492,420,429,490&search=)、[分子動力学](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519,8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519,8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519,8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519&search=)、[量子化学](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519,8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519&search=)、[物理学](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281,6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281,6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281,6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281,6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281&search=)、および HPC における[世界最高性能のアプリケーション](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=58,59,60,293,98,172,223,227,228,265,487,488,114,389,220,258,461&search=)の多くに対応しています。\n",
    "\n",
    "CUDA の知識があれば、自身が開発したアプリケーションを高速化できるようになります。GPU 対応アプリケーションは、CPU 専用アプリケーションよりもはるかに高速で、CPU のパフォーマンスでは不可能だったコンピューティングを可能にします。このラボでは、CUDA C/C++ を使用した GPU 対応アプリケーションのプログラミングの基礎を学習します。これにより、自身の CPU 専用アプリケーションを高速化して、パフォーマンスを高めたり、まったく新しいコンピューティング分野を切り開いたりできます。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 前提条件\n",
    "\n",
    "このラボを効果的に活用するには、次のスキルを習得している必要があります。\n",
    "\n",
    "- C で変数宣言、ループ、if/else ステートメントを使用できる。\n",
    "- C で関数を定義して呼び出す。\n",
    "- C で配列を割り当てる。\n",
    "\n",
    "CUDA に関する事前知識は必要ありません。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 目標\n",
    "\n",
    "このラボを完了すると、次のことができるようになります。\n",
    "\n",
    "- CPU 関数の呼び出しと GPU **カーネル起動**の両方を行う C/C++ プログラムを作成、コンパイル、実行する。\n",
    "- **実行構成**を使用して、並列**スレッド階層**を制御する。\n",
    "- シリアル ループをリファクタリングして、ループの反復を GPU 上で並列実行する。\n",
    "- CPU と GPU の両方で使用可能なメモリの割り当てと解放を行う。\n",
    "- CUDA コードによって生成されたエラーを処理する。\n",
    "- CPU 専用アプリケーションを高速化する。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Accelerated Systems\n",
    "\n",
    "*アクセラレーテッド システム*は、*異種混合システム*とも呼ばれるように、CPU と GPU の両方で構成されるシステムです。アクセラレーテッド システムは CPU プログラムを実行して、大規模並列処理を行う GPU の関数を起動します。このラボの環境は、NVIDIA GPU を搭載したアクセラレーテッド システムです。この GPU に関する情報は、nvidia-smi (*Systems Management Interface*) コマンド ライン コマンドを使用して照会できます。下のコード実行セルで Ctrl + Enter キーを押して、nvidia-smi コマンドを発行してください。このラボでコードを実行する場合には、必ず以下のようなセルで示されます。コマンドの出力結果は、コード実行セルのすぐ下に表示されます。すぐ下のコード実行ブロックを実行したら、出力結果で GPU の名前を確認してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvidia-smi"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## GPU アクセラレーテッド アプリケーションと CPU 専用アプリケーションの比較\n",
    "\n",
    "次のスライドは、このセクションの概要を視覚化したものです。内容を確認してから、次のセクションのトピック詳細に進んでください。\n",
    "\n",
    "<script>console.log('hi');</script>"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1R5BGqawxzawz35Qoge_bgfIRcLdw4qNr2bxbYDThgbk/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1R5BGqawxzawz35Qoge_bgfIRcLdw4qNr2bxbYDThgbk/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## GPU 用のアプリケーション コードの作成\n",
    "\n",
    "このラボで使用する C/C++ を始め、CUDA は一般的なプログラミング言語の拡張機能を広く提供しています。これにより、開発者はソース コード内の関数を GPU 上で簡単に実行することができます。\n",
    "\n",
    "以下は .cu (CUDA アクセラレーテッド プログラムの拡張子) ファイルです。ここには関数が 2 つ含まれています。最初の関数は CPU 上で、2 番目の関数は GPU 上で実行されます。定義方法と呼び出し方法について、各関数の違いをしっかりと確認してください。\n",
    "\n",
    "```cpp\n",
    "void CPUFunction()\n",
    "{\n",
    "  printf(\"This function is defined to run on the CPU.\\n\");\n",
    "}\n",
    "\n",
    "__global__ void GPUFunction()\n",
    "{\n",
    "  printf(\"This function is defined to run on the GPU.\\n\");\n",
    "}\n",
    "\n",
    "int main()\n",
    "{\n",
    "  CPUFunction();\n",
    "\n",
    "  GPUFunction<<<1, 1>>>();\n",
    "  cudaDeviceSynchronize();\n",
    "}\n",
    "```\n",
    "\n",
    "次に、重要なコード行と、アクセラレーテッド コンピューティングで使用される一般的な用語を紹介します。\n",
    "\n",
    "`__global__ void GPUFunction()`\n",
    "  - `__global__` キーワードは、その後の関数が GPU 上で実行され、**グローバルに**呼び出すことができることを示します (グローバルとは「CPU と GPU のいずれかで」という意味です)。\n",
    "  - 多くの場合、CPU 上で実行されるコードは**ホスト**コードと呼ばれ、GPU 上で実行されるコードは**デバイス**コードと呼ばれます。\n",
    "  - 戻り値の型 `void` に注目してください。`__global__` キーワードを使用して定義された関数は、`void` 型を返す必要があります。\n",
    "\n",
    "`GPUFunction<<<1, 1>>>();`\n",
    "  - 通常、GPU 上で実行する関数を呼び出す場合、この関数を**カーネル**と呼び、これを**起動**します。\n",
    "  - カーネルを起動するときは、**実行構成**を指定する必要があります。実行構成を指定するには、必要な引数をカーネルに渡す直前に `<<< ... >>>` 構文を使用します。\n",
    "  - 簡単に言うと、実行構成を使用することで、カーネルを起動するための**スレッド階層**を指定できます。スレッド階層では、スレッドのグループ (**ブロック**と呼ばれます) の数と、各ブロックで実行する**スレッド**の数を定義します。実行構成についてはラボの後半で詳しく説明しますが、当面は 1 つのスレッド (実行構成の 2 番目の引数) が含まれる 1 つのスレッド ブロック (実行構成の 1 番目の引数) でカーネルを起動します。\n",
    "\n",
    "`cudaDeviceSynchronize();`\n",
    "  - 多くの C/C++ コードとは異なり、カーネルの起動は**非同期**です。そのため、CPU コードは*カーネルの起動の完了を待たずに*処理を継続します。\n",
    "  - CUDA ランタイム内の `cudaDeviceSynchronize` 関数を呼び出すと、ホスト (CPU) コードはデバイス (GPU) コードが完了するまで待ってから、CPU 上での実行を再開します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: Hello GPU カーネルを作成する\n",
    "\n",
    "[`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/01-hello/01-hello-gpu.cu) (<---- リンクをクリックすると、ソース ファイルが別のタブで開き、編集できます) には、既に動作するプログラムが含まれています。2 つの関数が含まれており、どちらも「Hello from the CPU」というメッセージを出力します。目標は、ソース ファイル内の `helloGPU` 関数をリファクタリングして、GPU 上で実際に実行し、結果メッセージを出力することです。\n",
    "\n",
    "- アプリケーションをリファクタリングしたら、すぐ下の `nvcc` コマンドでコンパイルして実行してください (`Ctrl + Enter` キーでコード実行セルのコンテンツを実行できます)。[`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/01-hello/01-hello-gpu.cu) 内のコメントは作業の助けになります。行き詰まったときや、自分の作業を確認したい場合は、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/01-hello/solutions/01-hello-gpu-solution.cu)を参照してください。以下のコマンドでコンパイルして実行する前に、ファイルの変更を忘れずに保存してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o hello-gpu 01-hello/01-hello-gpu.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "[`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/01-hello/01-hello-gpu.cu) のリファクタリングが完了したら、以下の変更を加え、各変更の後に (`Ctrl + Enter` キーを押しながら上のコード実行セルをクリックして) コンパイルして実行してください。エラーが発生した場合は、エラー メッセージをよく読んでください。エラーを知っておくと、自身のコードのアクセラレーションを行うときに役立ちます。\n",
    "\n",
    "- カーネル定義から `__global__` キーワードを削除します。エラーの行番号に注目してください。このエラーの「configured」が何を意味するのかを考えてみてください。終わったら、`__global__` を元に戻します。\n",
    "- 実行構成を削除します。「configured」の意味が理解できたでしょうか。終わったら、実行構成を元に戻します。\n",
    "- `cudaDeviceSynchronize` の呼び出しを削除します。コードをコンパイルして実行する前に、結果を予測してみてください。カーネルは非同期に起動されること、および `cudaDeviceSynchronize` はカーネルの実行が完了するのを待ってからホストの実行を再開します。終わったら、`cudaDeviceSynchronize` を元に戻します。\n",
    "- `01-hello-gpu.cu` をリファクタリングして、`Hello from the CPU` の**前**に `Hello from the GPU` を出力させます。\n",
    "- `01-hello-gpu.cu` をリファクタリングして、`Hello from the GPU` を **2 回**出力させます (Hello from the CPU の**前**に 1 回、**後**に 1 回)。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### CUDA コードのコンパイルと実行\n",
    "\n",
    "このセクションでは、`.cu` プログラムをコンパイルして実行する、先ほどの `nvcc` コマンドについて詳しく説明します。\n",
    "\n",
    "CUDA プラットフォームには [**NVIDIA CUDA コンパイラ**](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html) (`nvcc`) が付属しており、CUDA アプリケーションのホスト及びデバイスコードをコンパイルできます。このラボでは、実際に使用する範囲で、`nvcc` について説明します。ラボの完了後、`nvcc` について詳しく知りたい方は、この[ドキュメント](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html)を参照してください。\n",
    "\n",
    "`nvcc` は、gcc の経験があるユーザーには非常に馴染み深いものです。たとえば、`some-CUDA.cu`` ファイルのコンパイルは次のコマンドで簡単に実行できます。\n",
    "\n",
    "`nvcc -arch=sm_70 -o out some-CUDA.cu -run`\n",
    "  - `nvcc` は、nvcc コンパイラを使用するためのコマンドです。 \n",
    "  - `some-CUDA.cu` は、コンパイルするファイルとして渡されます。\n",
    "  - `o` フラグは、コンパイルされたプログラムの出力ファイルを指定する際に使用します。\n",
    "  - `arch` フラグは、ファイルをコンパイルする**アーキテクチャ**を示します。この例では `sm_70` を指定しているので、このラボで使用している Volta GPU をターゲットにコンパイルされます。詳しくは、[`arch` フラグ](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#options-for-steering-gpu-code-generation)、 [仮想アーキテクチャ機能](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#virtual-architectures)、および [GPU 機能](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#gpu-feature-list)に関するドキュメントを参照してください。\n",
    "  - `run` フラグを指定すると、正常にコンパイルされたバイナリが実行されるので、便利です。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## CUDA のスレッド階層\n",
    "\n",
    "次のスライドは、このセクションの概要を視覚化したものです。内容を確認してから、次のセクションのトピック詳細に進んでください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1ksh6pEW0giIfjl4t6RFwLDxiKsH8RBuCOk4vKttNo_U/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1ksh6pEW0giIfjl4t6RFwLDxiKsH8RBuCOk4vKttNo_U/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## カーネルの並列起動\n",
    "\n",
    "実行構成では、複数の GPU **スレッド**で並列実行するカーネルの起動を詳細に指定できます。厳密に言うと、**スレッド ブロック**または単に**ブロック**と呼ばれるスレッドのグループの数と、各スレッド ブロックに含めるスレッドの数を指定できます。実行構成の構文は次のとおりです。\n",
    "\n",
    "`<<< NUMBER_OF_BLOCKS, NUMBER_OF_THREADS_PER_BLOCK>>>`\n",
    "\n",
    "**カーネル コードは、カーネルの起動時に構成されたすべてのスレッド ブロックのすべてのスレッドで実行されます。**\n",
    "\n",
    "したがって、`someKernel` という名前のカーネルが定義されている場合、以下のようになります。\n",
    "  - `someKernel<<<1, 1>>()` では、1 個のスレッドを含む 1 個のスレッド ブロックで実行するように構成されるので、1 回だけ実行されます。\n",
    "  - `someKernel<<<1, 10>>()` では、10 個のスレッドを含む 1 個のスレッド ブロックで実行するように構成されるので、10 回実行されます。\n",
    "  - `someKernel<<<10, 1>>()` では、それぞれ 1 個のスレッドを含む 10 個のスレッド ブロックで実行するように構成されるので、10 回実行されます。\n",
    "  - `someKernel<<<10, 10>>()` では、それぞれ 10 個のスレッドを含む 10 個のスレッド ブロックで実行するように構成されるので、100 回実行されます。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: カーネルを並列起動する\n",
    "\n",
    "[`01-first-parallel.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/02-first-parallel/01-basic-parallel.cu) は、`This should be running in parallel.` というメッセージを出力する基本的な関数を呼び出します。以下の手順に従って、最初に GPU 上で実行し、次に単一で実行し、その後で複数のスレッド ブロックで 並列実行するようにリファクタリングしてください。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/02-first-parallel/solutions/01-basic-parallel-solution.cu)を参照してください。\n",
    "\n",
    "- `firstParallel` 関数をリファクタリングして、GPU 上で CUDA カーネルとして起動させます。`01-basic-parallel.cu` をコンパイルして実行した後にすぐ `nvcc`` コマンドを実行することで関数の出力を確認できます。\n",
    "- `firstParallel` カーネルをリファクタリングして、5 つのスレッドで並列実行させ、すべてが 1 つのスレッド ブロックで実行されるようにします。コードをコンパイルして実行したら、メッセージが 5 回出力されることを確認してください。\n",
    "- `firstParallel` カーネルを再度リファクタリングします。今回は、それぞれ 5 つのスレッドを含む 5 つのスレッド ブロックで並列実行させます。コンパイルして実行したら、今度はメッセージが 25 回出力されることを確認してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o basic-parallel 02-first-parallel/01-basic-parallel.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "\n",
    "## CUDA が提供するスレッド階層変数\n",
    "\n",
    "次のスライドは、このセクションの概要を視覚化したものです。内容を確認してから、次のセクションで説明しているトピックの詳細に進んでください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/15lskPIgFpZWswMC6NmhWlbMTN7Z3nu-oWvydJuHpHyw/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/15lskPIgFpZWswMC6NmhWlbMTN7Z3nu-oWvydJuHpHyw/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## スレッド インデックスとブロック インデックス\n",
    "\n",
    "各スレッドには、スレッド ブロック内で 0 から始まるインデックスが割り当てられます。さらに、各ブロックにも 0 から始まるインデックスが割り当てられます。スレッドがスレッド ブロックにグループ化されるのと同様に、ブロックは**グリッド**にグループ化されます。グリッドは CUDA スレッド階層の中で最上位のエンティティです。つまり、CUDA カーネルは、1 つ以上のブロックから成るグリッドで実行され、各ブロックにはそれぞれ同じ数 (1 つ以上) のスレッドが含まれます。\n",
    "\n",
    "CUDA カーネルは、カーネルを実行している (ブロック内の) スレッドのインデックスと、スレッドが含まれている (グリッド内の) ブロックのインデックスの両方を返す特別な変数にアクセスできます。これが `threadIdx.x` 変数と `blockIdx.x` 変数です。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: 特定のスレッド インデックスとブロック インデックスを使用する\n",
    "\n",
    "[`01-thread-and-block-idx.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/03-indices/01-thread-and-block-idx.cu) ファイルには、失敗メッセージを出力するカーネルが含まれています。このファイルを開いて、成功メッセージを出力するように実行構成を更新する方法を学習します。リファクタリングが完了したら、下のコード実行セルでコードをコンパイルして実行し、作業を確認してください。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/03-indices/solutions/01-thread-and-block-idx-solution.cu)を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o thread-and-block-idx 03-indices/01-thread-and-block-idx.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## for ループの高速化\n",
    "\n",
    "CPU 専用アプリケーションの for ループは、高速化に適しています。ループの各反復を順番に実行する代わりに、ループの各反復をそれぞれのスレッドで並列実行することができます。以下の for ループを見てみましょう。ループの実行回数を定義して、ループの各繰り返し処理を制御していることがわかります。\n",
    "\n",
    "```cpp\n",
    "int N = 2<<20;\n",
    "for (int i = 0; i < N; ++i)\n",
    "{\n",
    "  printf(\"%d\\n\", i);\n",
    "}\n",
    "```\n",
    "\n",
    "このループを並列化するには、次の 2 つの手順を実行する必要があります。\n",
    "\n",
    "- **ループ 1 回分の処理**を行うカーネルを作成する必要があります。\n",
    "- カーネルは実行中の他のカーネルに依存しないので、カーネルを適切な数だけ実行するように実行構成 (ループの反復回数など) を指定します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: 1 つのスレッド ブロックで for ループを高速化する\n",
    "\n",
    "[`01-single-block-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/04-loops/01-single-block-loop.cu)  内の loop 関数は、`0` ～ `9` までの数字を連続して出力する for ループを実行します。この `loop` 関数を、起動後に `N` 回の反復を並行で実行する CUDA カーネルにリファクタリングしてください。リファクタリングが正常に終了した後も、`0` ～ `9` の数字が出力される必要があります。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/04-loops/solutions/01-single-block-loop-solution.cu) を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o single-block-loop 04-loops/01-single-block-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 並列スレッドの調整\n",
    "\n",
    "次のスライドは、このセクションの概要を視覚化したものです。内容を確認してから、次のセクションのトピック詳細に進んでください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1s2KpXDEX0OumWLKui6y7F4_IkqO8okIy1lE-lnzA48g/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1s2KpXDEX0OumWLKui6y7F4_IkqO8okIy1lE-lnzA48g/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## さらなる並列化のためのブロック次元の使用\n",
    "\n",
    "1 つのスレッド ブロックに存在できるスレッドの数は、最大で 1,024 です。アクセラレーテッド アプリケーションで並列処理数を増やすためには、複数のスレッド ブロックを連携できなければなりません。\n",
    "\n",
    "CUDA カーネルは、ブロック内のスレッド数を返す特別な変数 `blockDim.x` にアクセスできます。この変数を `blockIdx.x` と `threadIdx.x` を組み合わせた慣用表現 `threadIdx.x + blockIdx.x * blockDim.x` で使用して、複数のスレッドから成る複数のブロックで並列実行を編成することで、高度な並列化を実現できます。以下は、詳細に示した例です。\n",
    "\n",
    "実行構成 `<<<10, 10>>>`` では、10 個のスレッドから成る 10 個のブロックにより、合計 100 個のスレッドのグリッドが起動します。したがって、各スレッドで `0`` ～ `99`` の間の一意のインデックスを計算する必要があります。\n",
    "\n",
    "- ブロックの `blockIdx.x` が `0` の場合、`blockIdx.x * blockDim.x` は `0` になります。`0` を、`threadIdx.x`` の取り得る値 `0` ～ `9` に加算すると、100 個のスレッドのグリッド内で `0` ～ `9` のインデックスを生成できます。\n",
    "- ブロックの `blockIdx.x` が `1` の場合、`blockIdx.x * blockDim.x` は `10` になります。`10` を、`threadIdx.x` の取り得る値 `0` ～ `9` に加算すると、100 個のスレッドのグリッド内で `10` ～ `19` のインデックスを生成できます。\n",
    "- ブロックの `blockIdx.x` が `5` の場合、`blockIdx.x * blockDim.x` は `50` になります。`50` を、`threadIdx.x` の取り得る値 `0` ～ `9` に加算すると、100 個のスレッドのグリッド内で `50` ～ `59` のインデックスを生成できます。\n",
    "- ブロックの `blockIdx.x` が `9` の場合、`blockIdx.x` * `blockDim.x` は `90` になります。`90` を、`threadIdx.x` の取り得る値 `0` ～ `9` に加算すると、100 個のスレッドのグリッド内で `90` ～ `99` のインデックスを生成できます。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: 複数のスレッド ブロックで for ループを高速化する\n",
    "\n",
    "[`02-multi-block-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/04-loops/02-multi-block-loop.cu)内の `loop` 関数は、`0` ～ `9` までの数字を順に出力する `for` ループを実行します。この `loop` 関数を、起動すると `N` 回の反復を並行で実行する CUDA カーネルにリファクタリングしてください。リファクタリングが正常に終了した後も、`0` ～ `9` の数字が出力されます。この演習では、追加の制約として、*2 個以上のスレッド ブロック*を起動する実行構成を使用してください。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/04-loops/solutions/02-multi-block-loop-solution.cu) を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o multi-block-loop 04-loops/02-multi-block-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## GPU と CPU でアクセス可能なメモリの割り当て\n",
    "\n",
    "バージョン 6 以降の CUDA では、CPU ホストと任意の数の GPU デバイスのどちらにも簡単に利用可能なメモリを割り当てることができます。アクセラレーテッド アプリケーションで最適なパフォーマンスを引き出すための[中級や上級のメモリ管理手法](http://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#memory-optimizations)はたくさんありますが、ここで紹介する CUDA の基本的なメモリ管理手法では、手間をかけずに、CPU 専用アプリケーションよりも大幅にパフォーマンスを向上させることができます。\n",
    "\n",
    "メモリの割り当ておよび解放を行い、ホスト コードとデバイス コードの両方で参照可能なポインターを取得するには、次の例のように `malloc` と `free` の呼び出しを `cudaMallocManaged` と `cudaFree` に置き換えます。\n",
    "\n",
    "```cpp\n",
    "// CPU-only\n",
    "\n",
    "int N = 2<<20;\n",
    "size_t size = N * sizeof(int);\n",
    "\n",
    "int *a;\n",
    "a = (int *)malloc(size);\n",
    "\n",
    "// Use `a` in CPU-only program.\n",
    "\n",
    "free(a);\n",
    "```\n",
    "\n",
    "```cpp\n",
    "// Accelerated\n",
    "\n",
    "int N = 2<<20;\n",
    "size_t size = N * sizeof(int);\n",
    "\n",
    "int *a;\n",
    "// Note the address of `a` is passed as first argument.\n",
    "cudaMallocManaged(&a, size);\n",
    "\n",
    "// Use `a` on the CPU and/or on any GPU in the accelerated system.\n",
    "\n",
    "cudaFree(a);\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: ホストとデバイスの両方で配列を操作する\n",
    "\n",
    "The [`01-double-elements.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/05-allocate/01-double-elements.cu) のプログラムは、配列を割り当て、ホスト上で整数値に初期化し、GPU 上でこれらの各値を並列で 2 倍にした後、この処理が成功したかどうかをホスト上で確認します。現在、このプログラムは動作しません。ホストとデバイスの両方でポインター a の配列操作を試みますが、(malloc を使用して) 割り当てられた配列にホスト上でアクセスするだけです。次の条件を満たすようにアプリケーションをリファクタリングしてください。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/05-allocate/solutions/01-double-elements-solution.cu)を参照してください。\n",
    "\n",
    "- `a` を、ホスト コードとデバイス コードの両方で使用できるようにします。\n",
    "- `a` のメモリを適切に解放します。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o double-elements 05-allocate/01-double-elements.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## グリッド サイズと処理量の不均衡\n",
    "\n",
    "次のスライドは、このセクションの概要を視覚化したものです。内容を確認してから、次のセクションのトピック詳細に進んでください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1JAaCP73CsiSn5yeV3dmUg37vn2sOMYj0lBrrLKlwFWA/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1JAaCP73CsiSn5yeV3dmUg37vn2sOMYj0lBrrLKlwFWA/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## スレッド数とブロック構成の不均衡への対応\n",
    "\n",
    "ループの並列化に必要なスレッド数を生成するように実行構成を指定できないことがあります。\n",
    "\n",
    "よくある例として、最適なブロック サイズを選択したい場合が挙げられます。たとえば、GPU ハードウェアの特性上パフォーマンス向上に効果的なのが、ブロックのスレッド数を 32 の倍数にすることです。それぞれ 256 個のスレッド (32 の倍数) を含むブロックを起動し、1,000 個の並列タスクを実行する必要があるとします (わかりやすくするために処理数を減らしています)。1,000 は 32 の倍数 (整数値) ではないため、グリッド内で合計 1,000 個のスレッドを生成するブロック数は存在しません。\n",
    "\n",
    "このシナリオには、以下の方法で簡単に対処できます。\n",
    "\n",
    "- 割り当てられた処理に必要な数*以上*のスレッドの実行構成を作成します。\n",
    "- 処理するデータ セットの合計サイズを表す値 (つまり、処理を完了するために必要なスレッドの合計数) を引数 (`N`) としてカーネルに渡します。\n",
    "- (`tid+bid*bdim` を使用して) グリッド内のスレッドのインデックスを計算した後、このインデックスが `N` を超えていないかどうかを確認し、超えていない場合にのみカーネルの該当する処理を実行します。\n",
    "\n",
    "ここでは、`N` と 1 ブロック内のスレッドの数のどちらもわかっているが、グリッド内のスレッド数と `N` が完全に一致しない場合に実行構成を作成する慣用表現の例を示します。この方法では、少なくとも `N` に対して必要な数のスレッドが存在し、多くても余分に 1 ブロック分のスレッドしか追加されることはありません。\n",
    "\n",
    "```cpp\n",
    "// Assume `N` is known\n",
    "int N = 100000;\n",
    "\n",
    "// Assume we have a desire to set `threads_per_block` exactly to `256`\n",
    "size_t threads_per_block = 256;\n",
    "\n",
    "// Ensure there are at least `N` threads in the grid, but only 1 block's worth extra\n",
    "size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;\n",
    "\n",
    "some_kernel<<<number_of_blocks, threads_per_block>>>(N);\n",
    "```\n",
    "\n",
    "上記の実行構成では `N` よりも多くのスレッドがグリッド内に生成されるので、`some_kernel` が「余分な」スレッドで実行される場合に範囲外のデータ要素にアクセスしないように、次のように `some_kernel` の定義内で対処する必要があります。\n",
    "\n",
    "```cpp\n",
    "__global__ some_kernel(int N)\n",
    "{\n",
    "  int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
    "\n",
    "  if (idx < N) // Check to make sure `idx` maps to some value within `N`\n",
    "  {\n",
    "    // Only do work if it does\n",
    "  }\n",
    "}\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: 不均衡の実行構成で for ループを高速化する\n",
    "\n",
    "[`02-mismatched-config-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/05-allocate/02-mismatched-config-loop.cu) のプログラムは、整数値の 1,000 個の要素の配列に対して `cudaMallocManaged` を使用してメモリを割り当てた後、CUDA カーネルを使用して配列のすべての値を並列で初期化しようとします。このプログラムは、`N` と `threads_per_block`` の数の両方がわかっていることを前提としています。ここでのタスクは、次の 2 つの目標を達成することです。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/05-allocate/solutions/02-mismatched-config-loop-solution.cu)を参照してください。\n",
    "\n",
    "- 処理対象の `a` の要素と同じ数以上のスレッドを生成する値を `number_of_blocks` に代入します。\n",
    "- 範囲外のデータ要素を処理しないように `initializeElementsTo` カーネルを更新します。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o mismatched-config-loop 05-allocate/02-mismatched-config-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## グリッド ストライド ループ\n",
    "\n",
    "次のスライドは、このセクションの概要を視覚化したものです。内容を確認してから、次のセクションのトピック詳細に進んでください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1TAMVL5JJJ2A2_DcEAdqsbhq0r06LykLDgjz6GA3HyOw/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1TAMVL5JJJ2A2_DcEAdqsbhq0r06LykLDgjz6GA3HyOw/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## グリッドより大きいデータ セット\n",
    "\n",
    "実行パフォーマンスの高い実行構成を作成する場合がほとんどですが、やむを得ず、グリッド内のスレッドの数をデータ セットのサイズより少なくすることもあります。1,000 個の要素の配列と、250 個のスレッドのグリッドがあるとします (わかりやすく説明するためにサイズを小さくしています)。この場合、グリッド内の各スレッドを 4 回使用する必要があります。一般的な方法の 1 つが、カーネル内で**グリッド ストライド ループ**を使用することです。\n",
    "\n",
    "グリッド ストライド ループでは、各スレッドは `tid+bid*bdim` を使用してグリッド内で一意のインデックスを計算し、配列内のそのインデックスの要素に対して処理を実行した後、グリッド内のスレッド数をインデックスに加算し、インデックスが配列の範囲外になるまで繰り返します。たとえば、500 個の要素の配列と 250 個のスレッドのグリッドの場合、グリッド内のインデックス 20 のスレッドは次の処理を実行します。\n",
    "\n",
    "- 500 個の要素の配列の要素 20 に対して処理を実行します\n",
    "- このインデックスにグリッド サイズの 250 を加算します。そうすると、結果は 270 になります\n",
    "- 500 個の要素の配列の要素 270 に対して処理を実行します\n",
    "- このインデックスにグリッド サイズの 250 を加算します。そうすると、結果は 520 になります\n",
    "- 520 は配列の範囲外であるため、スレッドは処理を停止します\n",
    "\n",
    "CUDA には、グリッド内のブロック数を返す特別な変数 gridDim.x があります。グリッド内のスレッドの総数を計算するには、単純にグリッド内のブロック数に各ブロック内のスレッド数を掛けます (`gridDim.x * blockDim.x`)。これを利用したカーネル内のグリッド ストライド ループの詳細な例を次に示します。\n",
    "\n",
    "```cpp\n",
    "__global void kernel(int *a, int N)\n",
    "{\n",
    "  int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;\n",
    "  int gridStride = gridDim.x * blockDim.x;\n",
    "\n",
    "  for (int i = indexWithinTheGrid; i < N; i += gridStride)\n",
    "  {\n",
    "    // do work on a[i];\n",
    "  }\n",
    "}\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: グリッド ストライド ループを使用してグリッドよりも大きな配列を操作する\n",
    "\n",
    "[`03-grid-stride-double.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/05-allocate/03-grid-stride-double.cu) をリファクタリングして、`doubleElements` カーネルでグリッド ストライド ループを使用し、`N` よりも小さいグリッドでスレッドを再利用して、配列のすべての要素を処理できるようにしてください。このプログラムは、配列のすべての要素が 2 倍になっているかどうかを出力します。現在は正確に `FALSE` を出力します。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/05-allocate/solutions/03-grid-stride-double-solution.cu)を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o grid-stride-double 05-allocate/03-grid-stride-double.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## エラー処理\n",
    "\n",
    "他のアプリケーションと同様、 CUDA コードにもエラー処理は不可欠です。CUDA 関数の多くは ([メモリ管理関数](http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY)など)、`cudaError_t` 型の値を返します。この値を使用することで、関数の呼び出し中にエラーが発生したかどうかを確認できます。次に、`cudaMallocManaged` の呼び出しに対するエラー処理の例を示します。\n",
    "\n",
    "```cpp\n",
    "cudaError_t err;\n",
    "err = cudaMallocManaged(&a, N)                    // Assume the existence of `a` and `N`.\n",
    "\n",
    "if (err != cudaSuccess)                           // `cudaSuccess` is provided by CUDA.\n",
    "{\n",
    "  printf(\"Error: %s\\n\", cudaGetErrorString(err)); // `cudaGetErrorString` is provided by CUDA.\n",
    "}\n",
    "```\n",
    "\n",
    "`void` を返すように定義されているカーネルを起動した場合、`cudaError_t` 型の値が返されません。起動構成に誤りがある場合など、カーネル起動時のエラーを確認するために、CUDA には `cudaGetLastError` 関数が用意されています。この関数は `cudaError_t` 型の値を返します。\n",
    "\n",
    "```cpp\n",
    "/*\n",
    " * This launch should cause an error, but the kernel itself\n",
    " * cannot return it.\n",
    " */\n",
    "\n",
    "someKernel<<<1, -1>>>();  // -1 is not a valid number of threads.\n",
    "\n",
    "cudaError_t err;\n",
    "err = cudaGetLastError(); // `cudaGetLastError` will return the error from above.\n",
    "if (err != cudaSuccess)\n",
    "{\n",
    "  printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
    "}\n",
    "```\n",
    "\n",
    "非同期カーネルの実行中などに非同期に発生したエラーを検知するには、後続の同期 CUDA ランタイム API 呼び出しから返されるステータスを確認する必要があります。たとえば、`cudaDeviceSynchronize` は、事前に起動したカーネルが 1 つでも失敗するとエラーを返します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: エラー処理を追加する\n",
    "\n",
    "[`01-add-error-handling.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/06-errors/01-add-error-handling.cu) をコンパイルして実行すると、配列の要素が正常に 2 倍にならない場合に結果を出力します。しかし、プログラム内のエラーは示されません。CUDA エラーを処理するようにこのアプリケーションをリファクタリングして、プログラムの問題点を効果的にデバッグできるようにしてください。CUDA 関数の呼び出し時に発生しやすい同期エラーと、CUDA カーネルの実行中に発生しやすい非同期エラーの両方を調べる必要があります。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/06-errors/solutions/01-add-error-handling-solution.cu) を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o add-error-handling 06-errors/01-add-error-handling.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### CUDA のエラー処理関数\n",
    "\n",
    "エラー確認用の CUDA 関数呼び出しをラップするマクロを作成すると便利です。次に例を示します。このマクロは演習で自由にご利用ください。\n",
    "\n",
    "```cpp\n",
    "#include <stdio.h>\n",
    "#include <assert.h>\n",
    "\n",
    "inline cudaError_t checkCuda(cudaError_t result)\n",
    "{\n",
    "  if (result != cudaSuccess) {\n",
    "    fprintf(stderr, \"CUDA Runtime Error: %s\\n\", cudaGetErrorString(result));\n",
    "    assert(result == cudaSuccess);\n",
    "  }\n",
    "  return result;\n",
    "}\n",
    "\n",
    "int main()\n",
    "{\n",
    "\n",
    "/*\n",
    " * The macro can be wrapped around any function returning\n",
    " * a value of type `cudaError_t`.\n",
    " */\n",
    "\n",
    "  checkCuda( cudaDeviceSynchronize() )\n",
    "}\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## まとめ\n",
    "\n",
    "ここまでで、以下のラボの目標をすべて達成しました。\n",
    "\n",
    "- CPU 関数の呼び出しと GPU **カーネル**の**起動**の両方を行う C/C++ プログラムを作成、コンパイル、実行する。\n",
    "- **実行構成**を使用して、並列**スレッド階層**を制御する。Control parallel **thread hierarchy** using **execution configuration**.\n",
    "- シリアル ループをリファクタリングして、ループの反復を GPU 上で並列実行する。\n",
    "- CPU と GPU の両方で使用可能なメモリの割り当てと解放を行う。\n",
    "- CUDA コードによって生成されたエラーを処理する。\n",
    "\n",
    "次は、以下のラボの最終目標を達成します。\n",
    "\n",
    "- CPU 専用アプリケーションを高速化する。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 最後の演習: ベクトル加法アプリケーションを高速化する\n",
    "\n",
    "次の課題では、このラボでこれまでに学んだことをすべて活用します。CPU 専用のベクトル加法プログラムをアクセラレートします。これはあまり洗練されたプログラムではありませんが、CUDA を使用したアプリケーションの GPU アクセラレーションの知識を応用するのに適しています。この演習後に余裕があれば、ぜひ*上級者向けの課題* セクションに進んで、より複雑なコードベースを使用した課題に取り組んでください。\n",
    "\n",
    "[`01-vector-add.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/07-vector-add/01-vector-add.cu) には、動作する CPU 専用のベクトル加法アプリケーションが含まれています。その中の `addVectorsInto` 関数をアクセラレートして、GPU 上で CUDA カーネルとして実行し、並列で処理できるようにしてください。次のことを実現する必要があります。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/07-vector-add/solutions/01-vector-add-solution.cu)を参照してください。\n",
    "- `addVectorsInto` の定義を CUDA カーネルに拡張します。\n",
    "- `addVectorsInto` が CUDA カーネルとして起動できるように実行構成を選択します。\n",
    "- 3 つのベクトル a、b、result にホスト コードおよびデバイス コードからアクセスできるように、メモリの割り当てと解放を更新します。\n",
    "- `addVectorsInto` の本体をリファクタリングします。1 つのスレッド内で起動し、入力ベクトルに対して 1 スレッド分の処理のみを実行する必要があります。スレッドが入力ベクトルの範囲外の要素にアクセスしないように注意してください。また、スレッドで入力ベクトルの複数の要素を処理する必要があるかどうかを確認します。\n",
    "- エラーを返さずに失敗する可能性のある CUDA コードにエラー処理を追加します。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o vector-add 07-vector-add/01-vector-add.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 上級者向けの課題\n",
    "\n",
    "次の演習は、余裕のある方のために用意した追加の課題です。ここでは、スキャフォールディング機能も少なく、より高度な手法が必要になります。難易度は高いですが、スキルの向上に役立ちます。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 2 次元と 3 次元のグリッドとブロック\n",
    "\n",
    "グリッドとブロックは、最大 3 次元まで定義できます。複数の次元を定義しても、パフォーマンスに影響することはありませんが、2 次元行列などの複数次元のデータを処理する際に非常に役立ちます。2 次元または 3 次元のグリッドまたはブロックを定義するには、CUDA の `dim3` 型を次のように使用します。\n",
    "\n",
    "```cpp\n",
    "dim3 threads_per_block(16, 16, 1);\n",
    "dim3 number_of_blocks(16, 16, 1);\n",
    "someKernel<<<number_of_blocks, threads_per_block>>>();\n",
    "```\n",
    "\n",
    "上の例の場合、`someKernel` 内では、変数 `gridDim.x`、`gridDim.y`、`blockDim.x`、`blockDim.y` はすべて `16` になります。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: 2 次元行列乗算アプリケーションを高速化する\n",
    "\n",
    "[`01-matrix-multiply-2d.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/08-matrix-multiply/01-matrix-multiply-2d.cu) ファイルには、完全なホスト関数 `matrixMulCPU` が含まれています。ここでのタスクは、`matrixMulGPU` CUDA カーネルを構築することです。このソース コードは、2 つの関数で行列乗算を実行し、その結果を比較して、作成した CUDA カーネルの正確性を検証します。次のガイドラインに従って、作業してください。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/08-matrix-multiply/solutions/01-matrix-multiply-2d-solution.cu) を参照してください。\n",
    "\n",
    "- `x` と `y` の次元を `1` より大きい値に設定した 2 つの `dim3` 値を引数に指定した実行構成を作成する必要があります。\n",
    "- カーネルの内部では、通常どおりグリッド内に実行中のスレッドの一意のインデックスを確立しますが、スレッド 1 つに対してインデックスを 2 つ作成する必要があります。1 つはグリッドの `x` 軸用、もう 1 つは `y` 軸用です。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "FOUND ERROR at c[1][0]\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -arch=sm_70 -o matrix-multiply-2d 08-matrix-multiply/01-matrix-multiply-2d.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### 演習: 熱伝導率アプリケーションをアクセラレートする\n",
    "\n",
    "次の演習では、2 次元空間で銀の熱伝導をシミュレートするアプリケーションをアクセラレートします。\n",
    "\n",
    "[`01-heat-conduction.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/09-heat/01-heat-conduction.cu) 内の `step_kernel_mod` 関数を GPU 上で実行するように変換します。さらに、データを適切に割り当てて `main` 関数を変更し、CPU と GPU で使用できるようにしてください。`step_kernel_ref` 関数は CPU 上で実行され、エラー チェックに使用されます。このコードでは浮動小数点計算を行うため、プロセッサによって、また同じプロセッサでも単純な並べ替え操作によって、結果が若干異なる場合があります。そのため、このエラー チェック コードでは、完全一致ではなく、エラーのしきい値で判断します。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/01_AC_CUDA_C/09-heat/solutions/01-heat-conduction-solution.cu) を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o heat-conduction 09-heat/01-heat-conduction.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "> このタスクの元の熱伝導 CPU ソース コードのクレジットは、ヒューストン大学の [C ベースの熱伝導コードの OpenACC サンプル コード](http://docplayer.net/30411068-An-openacc-example-code-for-a-c-based-heat-conduction-code.html)に記載されています。"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 2",
   "language": "python",
   "name": "python2"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 2
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython2",
   "version": "2.7.12"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 1
}
